__attribute__((global, cooperative, reserve_r29_for_scavenge)) void kernel_upsample(int IC, int IH, int IW, int OH, int OW, float *__restrict dev_inp, float *__restrict dev_out)
{
    // 获取所有线程数量
    int thread_num = GetThreadNum();
    // 获取当前线程的全局索引
    int thread_idx = GetThreadIdx();

    // 定义与初始化DTE
    tops_dte_ctx_t dte_input[2];
    tops_dte_ctx_t dte_trans_input[2];
    tops_dte_ctx_t dte_trans_output[2];
    tops_dte_ctx_t dte_output[2];

    dte_input[0].init();
    dte_input[1].init();
    dte_trans_input[0].init();
    dte_trans_input[1].init();

    dte_trans_output[0].init();
    dte_trans_output[1].init();
    dte_output[0].init();
    dte_output[1].init();

    tops::event event_input0;
    tops::event event_input1;
    tops::event event_trans_input0;
    tops::event event_trans_input1;
    tops::event event_trans_output0;
    tops::event event_trans_output1;
    tops::event event_output0;
    tops::event event_output1;
    // 分块

    // 选取实际工作线程数, 在w维上并行
    int w_loops = OW;
    auto nr_working_threads = w_loops <= thread_num ? w_loops : thread_num;

    // 结束多余的工作线程
    if (thread_idx >= nr_working_threads)
    {
        return;
    }

    // 计算每个线程的工作量
    auto base_loops = w_loops / nr_working_threads;
    auto thread_w_loops = base_loops + ((base_loops * nr_working_threads + thread_idx) < w_loops ? 1 : 0);
    auto thread_h_loops = OH;
    auto thread_c_loops = IC;
    auto k_tile_size = tops::vlength<vfloat>() < IC ? tops::vlength<vfloat>() : IC;

    // 全局内存的形状
    int global_shape[3] = {IC, IH, IW};
    int global_output_shape[3] = {IC, OH, OW};

    // 私有内存的形状
    int private_shape[3] = {IC, 2, 2};
    int private_trans_shape[3] = {2, 2, IC};
    // int private_outtrans_shape[3] = {1, OH, IC};
    int private_output_shape[3] = {IC, OH, 1};
    auto L1_BUFFER_SIZE = IC * 2 * 2;
    auto L1_BUFFER_SIZE_OUT = 2048;

    // 定义与绑定内存
    tops::mdspan global_input(tops::Global, dev_inp, global_shape);
    tops::mdspan global_output(tops::Global, dev_out, global_output_shape);

    __valigned__ float l1_buffer_input0[L1_BUFFER_SIZE];
    __valigned__ float l1_buffer_input1[L1_BUFFER_SIZE];
    __valigned__ float l1_trans_buffer_input0[L1_BUFFER_SIZE];
    __valigned__ float l1_trans_buffer_input1[L1_BUFFER_SIZE];
    __valigned__ float l1_trans_buffer_output0[L1_BUFFER_SIZE_OUT];
    __valigned__ float l1_trans_buffer_output1[L1_BUFFER_SIZE_OUT];
    // __valigned__ float l1_buffer_output0[L1_BUFFER_SIZE_OUT];
    // __valigned__ float l1_buffer_output1[L1_BUFFER_SIZE_OUT];

    tops::mdspan l1_input0(tops::Private, l1_buffer_input0, private_shape);
    tops::mdspan l1_input1(tops::Private, l1_buffer_input1, private_shape);
    tops::mdspan l1_trans_input0(tops::Private, l1_trans_buffer_input0, private_trans_shape);
    tops::mdspan l1_trans_input1(tops::Private, l1_trans_buffer_input1, private_trans_shape);
    // tops::mdspan l1_trans_output0(tops::Private, l1_trans_buffer_output0, private_outtrans_shape);
    // tops::mdspan l1_trans_output1(tops::Private, l1_trans_buffer_output1, private_outtrans_shape);
    // tops::mdspan l1_output0(tops::Private, l1_buffer_output0, private_output_shape);
    // tops::mdspan l1_output1(tops::Private, l1_buffer_output1, private_output_shape);

    // DTE绑定
    dte_input[0].connect(dte_trans_input[0]);
    dte_input[1].connect(dte_trans_input[1]);
    dte_trans_output[0].connect(dte_output[0]);
    dte_trans_output[1].connect(dte_output[1]);

    // 采样过程
    for (auto i = 0; i < thread_w_loops; i++) // w方向并行
    {
        // 计算维度w的全局索引
        auto global_w_idx = thread_idx + i * nr_working_threads;
        for (auto j = 0; j < thread_h_loops; j++) // H方向
        {
            auto global_h_idx = j;

            // 映射到输入空间
            float iw_f = global_w_idx * (IW - 1) / (float)(OW - 1);
            float ih_f = global_h_idx * (IH - 1) / (float)(OH - 1);

            // 确定输入空间的四个邻近像素
            int ih0 = floorf(ih_f);
            int iw0 = floorf(iw_f);
            int ih1 = ih0 + 1 < IH ? ih0 + 1 : ih0;
            int iw1 = iw0 + 1 < IW ? iw0 + 1 : iw0;

            // 计算小数部分
            float delta_h = ih_f - ih0;
            float delta_w = iw_f - iw0;

            // 计算双线性插值的权重
            float w1 = (1.0f - delta_h) * (1.0f - delta_w);
            float w2 = (1.0f - delta_h) * delta_w;
            float w3 = delta_h * (1.0f - delta_w);
            float w4 = delta_h * delta_w;

            // 加载输入数据
            event_input0 = tops::slice_async(dte_input[0], l1_input0, global_input, {0, ih0, iw0});
            event_trans_input0 = tops::transpose_async(dte_trans_input[0], l1_trans_input0, l1_input0, {1, 2, 0});

            tops::wait(event_trans_input0);

        
            for (auto k = 0; k < thread_c_loops; k += k_tile_size)
            {
                const auto &val1 = tops::vload<vfloat>(l1_trans_buffer_input0 + k);
                const auto &val2 = tops::vload<vfloat>(l1_trans_buffer_input0 + k + IC);
                const auto &val3 = tops::vload<vfloat>(l1_trans_buffer_input0 + k + IC * 2);
                const auto &val4 = tops::vload<vfloat>(l1_trans_buffer_input0 + k + IC * 3);

                // 与权重相乘
                // 广播权重
                const auto &v_w1 = tops::vbroadcast<vfloat>(w1);
                const auto &v_w2 = tops::vbroadcast<vfloat>(w2);
                const auto &v_w3 = tops::vbroadcast<vfloat>(w3);
                const auto &v_w4 = tops::vbroadcast<vfloat>(w4);

                const auto &val_w1 = tops::vmul<vfloat>(val1, v_w1);
                const auto &val_w2 = tops::vmul<vfloat>(val2, v_w2);
                const auto &val_w3 = tops::vmul<vfloat>(val3, v_w3);
                const auto &val_w4 = tops::vmul<vfloat>(val4, v_w4);

                // 加权求和
                const auto &val_sum1 = tops::vadd<vfloat>(val_w1, val_w2);
                const auto &val_sum2 = tops::vadd<vfloat>(val_w3, val_w4);
                const auto &val_sum = tops::vadd<vfloat>(val_sum1, val_sum2);

                // 写入输出
                tops::vstore(val_sum, l1_trans_buffer_output0 + global_h_idx * IC + k);
            }
        }

        event_trans_output0 = tops::transpose_async(dte_trans_output[0], l1_output0, l1_trans_output0, {2, 1, 0});
        event_output0 = tops::deslice_async(dte_output[0], global_output, l1_output0, {0, 0, global_w_idx});
        tops::wait(event_output0);
    }

    // 销毁DTE
    dte_input[0].destroy();
    dte_input[1].destroy();
    dte_trans_input[0].destroy();
    dte_trans_input[1].destroy();
    dte_trans_output[0].destroy();
    dte_trans_output[1].destroy();
    dte_output[0].destroy();
    dte_output[1].destroy();
}
void GCU_UpSample(int IC, int IH, int IW, int OH, int OW, float *__restrict dev_inp, float *__restrict dev_out)
{
    // 设置kernel参数
    static const size_t blocks = 2;
    static const size_t threads = 12;

    // 调用kernel
    kernel_upsample<<<blocks, threads>>>(IC, IH, IW, OH, OW, dev_inp, dev_out);
}